#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;    
template<typename TYPE_X, typename TYPE_Y> class KernelDawsn {
    using T = TYPE_X;
public:
    __aicore__ inline KernelDawsn() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
                                uint32_t CoreDataNum, uint32_t finalTileNum, uint32_t tileDataNum, uint32_t TailDataNum, int32_t ySize) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

        this->coreDataNum = CoreDataNum;
        this->tileNum = finalTileNum;
        this->tileDataNum = tileDataNum;
        this->tailDataNum = TailDataNum;
        this->ySize = ySize;

        xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x, this->coreDataNum);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y, this->coreDataNum);

        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_X));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(DTYPE_Y));

        pipe.InitBuffer(QueueTmp1, 32 * sizeof(float));
        pipe.InitBuffer(QueueTmp2, 32 * sizeof(float));
        pipe.InitBuffer(QueueTmp3, 32 * sizeof(float));
        pipe.InitBuffer(QueueTmp4, 32 * sizeof(float));
        pipe.InitBuffer(QueueTmp5, 32 * sizeof(float));
    }
    __aicore__ inline void Process() {
    float COEF_AN[] = {1.13681498971755972054E-11,
               8.49262267667473811108E-10,
               1.94434204175553054283E-8,
               9.53151741254484363489E-7,
               3.07828309874913200438E-6,
               3.52513368520288738649E-4,
               -8.50149846724410912031E-4,
               4.22618223005546594270E-2,
               -9.17480371773452345351E-2,
               9.99999999999999994612E-1};
    int COEF_AN_COUNT = 9;
    float COEF_AD[] = {2.40372073066762605484E-11,
               1.48864681368493396752E-9,
               5.21265281010541664570E-8,
               1.27258478273186970203E-6,
               2.32490249820789513991E-5,
               3.25524741826057911661E-4,
               3.48805814657162590916E-3,
               2.79448531198828973716E-2,
               1.58874241960120565368E-1,
               5.74918629489320327824E-1,
               1.00000000000000000539E0};
    int COEF_AD_COUNT = 10;

    float COEF_BN[] = {5.08955156417900903354E-1,
               -2.44754418142697847934E-1,
               9.41512335303534411857E-2,
               -2.18711255142039025206E-2,
               3.66207612329569181322E-3,
               -4.23209114460388756528E-4,
               3.59641304793896631888E-5,
               -2.14640351719968974225E-6,
               9.10010780076391431042E-8,
               -2.40274520828250956942E-9,
               3.59233385440928410398E-11};
    int COEF_BN_COUNT = 10;
    float COEF_BD[] = {-6.31839869873368190192E-1,
               2.36706788228248691528E-1,
               -5.31806367003223277662E-2,
               8.48041718586295374409E-3,
               -9.47996768486665330168E-4,
               7.81025592944552338085E-5,
               -4.55875153252442634831E-6,
               1.89100358111421846170E-7,
               -4.91324691331920606875E-9,
               7.18466403235734541950E-11};
    int COEF_BD_COUNT = 10;

    float COEF_CN[] = {-5.90592860534773254987E-1,
               6.29235242724368800674E-1,
               -1.72858975380388136411E-1,
               1.64837047825189632310E-2,
               -4.86827613020462700845E-4};
    int COEF_CN_COUNT = 4;
    float COEF_CD[] = {-2.69820057197544900361E0,
               1.73270799045947845857E0,
               -3.93708582281939493482E-1,
               3.44278924041233391079E-2,
               -9.73655226040941223894E-4};
    int COEF_CD_COUNT = 5;

    float THRESHOLD_3_25 = 3.25;
    float THRESHOLD_6_25 = 6.25;
    float THRESHOLD_1E_9 = 1.0e9;
        DTYPE_X sign;

        auto tmpx = QueueTmp1.Get<DTYPE_X>();
        auto tmpy = QueueTmp2.Get<DTYPE_X>();
        auto tmpdata_square = QueueTmp3.Get<DTYPE_X>();
        auto tmpdata_polevl_an = QueueTmp4.Get<DTYPE_X>();
        auto tmpdata_polevl_ad = QueueTmp5.Get<DTYPE_X>();

        tmpdata_polevl_ad.SetValue(0, (DTYPE_X)COEF_AN[0]);
        for (int32_t i = 0; i < this->ySize; i++)
        {
            float x = xGm.GetValue(i);
            float y;
            sign = 1;
            if((float)x < 0)
            {
                x = -(float)x;
                sign = -1;
            }

            tmpx.SetValue(0, (float)x);

            if((float)x < (float)3.25)
            {
                float data_square = (float)x*(float)x;             
                float data_polevl_an = COEF_AN[0];
                for(int32_t index=1; index<=COEF_AN_COUNT; index++)
                {
                    data_polevl_an = (float)data_polevl_an*(float)data_square;
                    data_polevl_an = (float)data_polevl_an + (float)COEF_AN[index];
                }

                float data_polevl_ad = (float)COEF_AD[0];
                for(int32_t index=1; index<=COEF_AD_COUNT; index++)
                {
                    data_polevl_ad = (float)data_polevl_ad*(float)data_square;
                    data_polevl_ad = (float)data_polevl_ad + (float)COEF_AD[index];
                }

                y = (float)x * (float)data_polevl_an;
                y = (float)y / (float)data_polevl_ad;
            }
            else if((float)x < (float)6.25)
            {
                float data_temp = (float)x*(float)x;
                data_temp = 1/data_temp;
                float data_rec = 1/x;
                
                float data_polevl_bn = COEF_BN[0];
                for(int32_t index=1; index<=COEF_BN_COUNT; index++)
                {
                    data_polevl_bn = (float)data_polevl_bn*(float)data_temp;
                    data_polevl_bn = (float)data_polevl_bn + (float)COEF_BN[index];
                }
                data_polevl_bn = data_polevl_bn * data_temp;
                float data_plevl_bd = (float)COEF_BD[0];
                data_plevl_bd = data_plevl_bd + data_temp;
                for(int32_t index=1; index<COEF_BD_COUNT; index++)
                {
                    data_plevl_bd = (float)data_plevl_bd*(float)data_temp;
                    data_plevl_bd = (float)data_plevl_bd + (float)COEF_BD[index];
                }
                data_plevl_bd = data_plevl_bd * x;
                y = data_polevl_bn/data_plevl_bd;
                y = data_rec + y;
                y = (float)0.5*y;
            }
            else if((float)x <= (float)1.0e9)
            {
                float data_temp = (float)x*(float)x;
                data_temp = 1/data_temp;
                float data_rec = 1/x;
                
                float data_polevl_cn = COEF_CN[0];
                for(int32_t index=1; index<=COEF_CN_COUNT; index++)
                {
                    data_polevl_cn = (float)data_polevl_cn*(float)data_temp;
                    data_polevl_cn = (float)data_polevl_cn + (float)COEF_CN[index];
                }
                data_polevl_cn = data_polevl_cn * data_temp;
                float data_plevl_cd = (float)COEF_CD[0];
                data_plevl_cd = data_plevl_cd + data_temp;
                for(int32_t index=1; index<COEF_CD_COUNT; index++)
                {
                    data_plevl_cd = (float)data_plevl_cd*(float)data_temp;
                    data_plevl_cd = (float)data_plevl_cd + (float)COEF_CD[index];
                }
                data_plevl_cd = data_plevl_cd * x;
                y = data_polevl_cn/data_plevl_cd;
                y = data_rec + y;
                y = (float)0.5*y;
            }
            else
            {
                y = (float)0.5/x;
            }
            y = (float)y * (float)sign;
            yGm.SetValue(i, (DTYPE_Y)y);
        }
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    TBuf<QuePosition::VECCALC> QueueTmp1, QueueTmp2,QueueTmp3,QueueTmp4, QueueTmp5;

    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_Y> yGm;
    uint32_t coreDataNum;
    uint32_t tileNum;
    uint32_t tileDataNum;
    uint32_t tailDataNum;
    uint32_t processDataNum;
    int32_t ySize;
};
extern "C" __global__ __aicore__ void dawsn(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    // TODO: user kernel impl
    KernelDawsn<DTYPE_X, DTYPE_Y> op;
    op.Init(x, y, 
            tiling_data.CoreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.TailDataNum, tiling_data.ySize);  
    op.Process();
}